<?xml version="1.0" encoding="ISO-8859-1"?>
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
  <meta http-equiv="Content-Type" content="text/html; charset=utf-8" />
<!-- blur.qdoc -->
  <title>Gaussian Blur Example</title>
  <link rel="stylesheet" type="text/css" href="style/style.css" />
  <script src="scripts/jquery.js" type="text/javascript"></script>
  <script src="scripts/functions.js" type="text/javascript"></script>
  <script src="./scripts/superfish.js" type="text/javascript"></script>
  <link rel="stylesheet" type="text/css" href="style/superfish.css" />  <script src="./scripts/narrow.js" type="text/javascript"></script>
  <link rel="stylesheet" type="text/css" href="style/narrow.css" />
  <!--[if IE]>
<meta name="MSSmartTagsPreventParsing" content="true">
<meta http-equiv="imagetoolbar" content="no">
<![endif]-->
<!--[if lt IE 7]>
<link rel="stylesheet" type="text/css" href="style/style_ie6.css">
<![endif]-->
<!--[if IE 7]>
<link rel="stylesheet" type="text/css" href="style/style_ie7.css">
<![endif]-->
<!--[if IE 8]>
<link rel="stylesheet" type="text/css" href="style/style_ie8.css">
<![endif]-->
</head>
<body class="" onload="CheckEmptyAndLoadList();">
<table border="0" cellpadding="0" cellspacing="0" width="100%">
<tr>
<td align="left" valign="top" width="150"><img src="images/qt-logo.png" align="left" border="0" />&nbsp;&nbsp;<img src="images/qtopencl-logo.png" border="0" /></td>
<td width="1">&nbsp;&nbsp;</td><td class="postheader" valign="center"> <a href="index.html"><font color="#004faf">Home</font></a>&nbsp;&middot; <a href="qtopencl-all-classes.html"><font color="#004faf">All Classes</font></a></td>
</table>              <li><a href="all-examples.html">Examples</a></li>              <li>Gaussian Blur Example</li><h1 class="title">Gaussian Blur Example</h1>
<span class="subtitle"></span>
<!-- $$$blur-description -->
<div class="descr"> <a name="details"></a>
<p>Files:</p>
<ul>
<li><a href="blur-blur-cl.html">blur/blur.cl</a></li>
<li><a href="blur-blur-qrc.html">blur/blur.qrc</a></li>
<li><a href="blur-blurtable-h.html">blur/blurtable.h</a></li>
<li><a href="blur-blurwidget-cpp.html">blur/blurwidget.cpp</a></li>
<li><a href="blur-blurwidget-h.html">blur/blurwidget.h</a></li>
<li><a href="blur-main-cpp.html">blur/main.cpp</a></li>
<li><a href="blur-blur-pro.html">blur/blur.pro</a></li>
<li><a href="blur-blur-qrc.html">blur/blur.qrc</a></li>
</ul>
<p>The Guassian Blur example shows how to blur QImage objects using OpenCL. The example animates blurring a 128x128 image from radius 1 to 16 and back again.</p>
<p class="centerAlign"><img src="images/blur-screenshot.png" /></p><p>We start by initializing a <a href="qclcontext.html">QCLContext</a> and then building the Gaussian blur program from the contents of the <tt>blur.cl</tt> resource file:</p>
<pre class="highlightedCode brush: cpp">        if (!context.create())
            qFatal(&quot;Could not create OpenCL context&quot;);

        program = context.buildProgramFromSourceFile(QLatin1String(&quot;:/blur.cl&quot;));</pre>
<p>Next, we load the Qt logo image and upload it into the GPU as <tt>srcImageBuffer</tt>:</p>
<pre class="highlightedCode brush: cpp">        QImage img(QLatin1String(&quot;:/qtlogo.png&quot;));
        srcImageBuffer = context.createImage2DCopy(img, QCLMemoryObject::ReadOnly);</pre>
<p><a href="qclcontext.html#createImage2DCopy">QCLContext::createImage2DCopy</a>() creates a <a href="qclimage2d.html">QCLImage2D</a> object that represents the image in the GPU. Here we have used <a href="qclmemoryobject.html#Access-enum">QCLMemoryObject::ReadOnly</a> to indicate that the GPU (not the CPU) will be reading from the image. The access modes are always with respect to what the GPU will be doing with them.</p>
<p>We will be performing the blur in two passes: the first pass processes the image horizontally and writes out a temporary image into <tt>tmpImageBuffer</tt>, and then the second pass processes the temporary image vertically and writes out the final image into <tt>dstImageBuffer</tt>:</p>
<pre class="highlightedCode brush: cpp">        tmpImageBuffer = context.createImage2DDevice(QImage::Format_ARGB32, img.size(), QCLMemoryObject::ReadWrite);

        dstImage = QImage(img.size(), QImage::Format_ARGB32);
        dstImageBuffer = context.createImage2DDevice(dstImage.format(), dstImage.size(), QCLMemoryObject::WriteOnly);</pre>
<p>This time we have used <a href="qclcontext.html#createImage2DDevice">QCLContext::createImage2DDevice</a>() to allocate memory in the GPU device. This memory is not directly accessible to the host CPU, which means that the GPU can allocate it within the most efficient memory area it can find. If we wanted the host to be able to directly access the memory, we would use <a href="qclcontext.html#createImage2DHost">QCLContext::createImage2DHost</a>() instead.</p>
<p>The temporary image buffer has an access mode of <a href="qclmemoryobject.html#Access-enum">QCLMemoryObject::ReadWrite</a> because it will be written by the horizontal pass and then read by the vertical pass. The destination image buffer has an access mode of <a href="qclmemoryobject.html#Access-enum">QCLMemoryObject::WriteOnly</a> because it is only written, never read, by the GPU.</p>
<p>We need two more buffers to hold the Gaussian weights and offsets for the specific blur radius:</p>
<pre class="highlightedCode brush: cpp">        weightsBuffer = context.createVector&lt;float&gt;(100);
        offsetsBuffer = context.createVector&lt;float&gt;(100);</pre>
<p>The last step of the initialization phase is to create the two kernels for the horizontal and vertical passes, set the OpenCL global work sizes to the size of the image, and set the local work sizes to the best sizes for 2D image processing:</p>
<pre class="highlightedCode brush: cpp">        hgaussian = program.createKernel(&quot;hgaussian&quot;);
        hgaussian.setGlobalWorkSize(img.size());
        hgaussian.setLocalWorkSize(hgaussian.bestLocalWorkSizeImage2D());

        vgaussian = program.createKernel(&quot;vgaussian&quot;);
        vgaussian.setGlobalWorkSize(img.size());
        vgaussian.setLocalWorkSize(vgaussian.bestLocalWorkSizeImage2D());</pre>
<p>Local work sizes break the image processing problem up into smaller chunks, usually 8x8 in size. These chunks can be processed more efficiently than linearly scanning the rows of the image in order.</p>
<p>When the program receives a paint event, we upload the precalculated Gaussian weights and offsets for the blur radius to the GPU:</p>
<pre class="highlightedCode brush: cpp">        offsetsBuffer.write(blurOffsets[radius], blurSizes[radius]);
        weightsBuffer.write(blurWeights[radius], blurSizes[radius]);</pre>
<p>Now we run the actual OpenCL kernel programs for the horizontal and vertical passes, passing the above OpenCL buffers as arguments:</p>
<pre class="highlightedCode brush: cpp">        hgaussian(srcImageBuffer, tmpImageBuffer, weightsBuffer, offsetsBuffer, blurSizes[radius]);
        vgaussian(tmpImageBuffer, dstImageBuffer, weightsBuffer, offsetsBuffer, blurSizes[radius]);</pre>
<p>Both functions have the following prototype:</p>
<pre class="highlightedCode brush: cpp">    __kernel void hgaussian(__read_only image2d_t srcImage,
                            __write_only image2d_t dstImage,
                            __global __read_only float *weights,
                            __global __read_only float *offsets,
                            const int numWeights)</pre>
<p>Once the kernels have been executed, we read the destination image out of the GPU and draw it as follows:</p>
<pre class="highlightedCode brush: cpp">        dstImageBuffer.read(&amp;dstImage);
        QPainter painter(this);
        painter.drawImage(0, 0, dstImage);</pre>
<p><a href="index.html">Return to Home</a></p>
</div>
<!-- @@@blur -->
<p /><address><hr /><div align="center">
<table width="100%" cellspacing="0" border="0"><tr class="address">
<td align="left">Copyright &copy; 2010 Nokia Corporation</td>
<td align="right">QtOpenCL Documentation</td>
</tr></table></div></address>  <script src="scripts/functions.js" type="text/javascript"></script>
  <script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-4457116-5']);
  _gaq.push(['_trackPageview']);
  (function() {
  var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
  ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
  var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
  </script>
</body>
</html>
